#define SDH_CALL static __device__ __host__

SDH_CALL float3 f3() { float3 r = {0.0f, 0.0f, 0.0f}; return r; }
SDH_CALL float4 f4() { float4 r = {0.0f, 0.0f, 0.0f, 0.0f}; return r; }

__global__ void kernel_00()
{
}

__global__ void kernel_01(float *pData)
{
	pData[threadIdx.x] = 0.0f;
}

__global__ void kernel_02(float3 * pData)
{
	pData[threadIdx.x] = f3();
}

__global__ void kernel_03(float4 *pData)
{
	pData[threadIdx.x] = f4();
}

__global__ void kernel_04(float *pData)
{
	pData[threadIdx.x*2] = 0.0f;
}

__global__ void kernel_05(float *pData)
{
	pData[threadIdx.x] = sinf((float)threadIdx.x);
}

__global__ void kernel_06(float *pData)
{
	pData[threadIdx.x] = __sinf((float)threadIdx.x);
}

__global__ void kernel_07(float *pData)
{
	pData[threadIdx.x] = cosf((float)threadIdx.x);
}

__global__ void kernel_08(float *pData)
{
	pData[threadIdx.x] = __cosf((float)threadIdx.x);
}

__global__ void kernel_09(float *pData)
{
	pData[threadIdx.x] = expf((float)threadIdx.x);
}

__global__ void kernel_10(float *pData)
{
	pData[threadIdx.x] = __expf((float)threadIdx.x);
}

__global__ void kernel_11(float4 *pDst, int * pSrc)
{
	float4 r;
	int a = pSrc[threadIdx.x];
	((float*)(&r))[a] = 0.0f;
	pDst[threadIdx.x] = r;
}

__global__ void kernel_12(float4 *pDst, int * pSrc)
{
	float4 r = {1.0f, 1.0f, 1.0f, 1.0f};
	int a = pSrc[threadIdx.x];
	r.x = (a==0 ? 0.0f: r.x);
	r.y = (a==1 ? 0.0f: r.y);
	r.z = (a==2 ? 0.0f: r.z);
	r.w = (a==3 ? 0.0f: r.w);

	pDst[threadIdx.x] = r;
}

#define DIM 128
#define V_X   smem[threadIdx.x+DIM*0]
#define V_Y   smem[threadIdx.x+DIM*1]
#define V_Z   smem[threadIdx.x+DIM*2]
#define V_W   smem[threadIdx.x+DIM*3]
#define V_(i) smem[threadIdx.x+DIM*(i)]
#define U_V(u)   u.x = V_X; u.y = V_Y; u.z = V_Z; u.w = V_W;

__global__ void kernel_13(float4 *pDst, int * pSrc)
{
	__shared__ float smem[DIM*4];
	float4 r = {1.0f, 1.0f, 1.0f, 1.0f};
	int a = pSrc[threadIdx.x];
	V_(a) = 0.0f;

	U_V(r);

	pDst[threadIdx.x] = r;
}


int main(int argc, char * argv[])
{
	kernel_00<<<dim3(1), dim3(512)>>>();

	kernel_01<<<dim3(1), dim3(512)>>>(NULL);

	kernel_02<<<dim3(1), dim3(512)>>>(NULL);

	kernel_03<<<dim3(1), dim3(512)>>>(NULL);

	return 0;
}